
#define MAX_TRIES 1000

#define NOT_PRESENT 0

#define HASH0(from, to, modulo) \
(to) = (from);           \
(to) -= ((to)<<6);    \
(to) ^= ((to)>>17);   \
(to) -= ((to)<<9);    \
(to) ^= ((to)<<4);    \
(to) -= ((to)<<3);    \
(to) ^= ((to)<<10);   \
(to) ^= ((to)>>15);   \
(to) %= (modulo);     \
(to) <<= 1;

#define HASH1(from, to, modulo)             \
(to) = (from);                              \
(to) ^= ((to) >> 20) ^ ((to) >> 12);        \
(to) = (to) ^ ((to) >> 7) ^ ((to) >> 4);    \
(to) %= (modulo);                           \
(to) <<= 1;

__kernel void cockooHash(__global uint * v, __global uint * hashTable, uint size ){
	int gid = get_global_id(0);
        //uint size = get_global_work_size();
        // unsigned int atomic_xchg ( 	volatile __global unsigned int *p unsigned int val)
        uint i = *(v + gid);

        for( int j = 0; j < MAX_TRIES; ++j){
            uint pos, newI;

            HASH0( i, pos, size );
            newI = atomic_xchg(hashTable + pos, i);
            // hashTable[pos];
            //hashTable[pos] = i;
            if (newI == NOT_PRESENT)
                break;
            i = newI;

            HASH1( i, pos, size );
            pos = pos + 1;
            newI = atomic_xchg(hashTable + pos, i);
            //newI = hashTable[pos];
            //hashTable[pos] = i;
            if (newI == NOT_PRESENT)
                break;
            i = newI;
        }
}
